#define K_MAX_SHAPE_DIM 0
#include "kernel_operator.h"
constexpr int32_t BUFFER_NUM = 2;
using namespace AscendC;
template <class T, class ComputeStrategy>
class KernelFresnelCos
{
    static constexpr float THRESHOLD_A = 2.5625f;
    static constexpr float Sqrt_A = 1.600781059358212171f;
    static constexpr float THRESHOLD_B = 36974.0f;
    static constexpr float THRESHOLD_PI = 3.1415926535897932384626433832795f;

    static constexpr float ReciprocalPI = 0.318309886183790671f;

    static constexpr float xishu = 1.5707963267948966f;
    // static constexpr float xishu2 = 0.636619772367581343f;
    static constexpr uint32_t SN_COUNT = 5;
    static constexpr uint32_t SD_COUNT = 6;
    static constexpr uint32_t CN_COUNT = 5;
    static constexpr uint32_t CD_COUNT = 6;
    static constexpr uint32_t FN_COUNT = 9;
    static constexpr uint32_t FD_COUNT = 10;
    static constexpr uint32_t GN_COUNT = 10;
    static constexpr uint32_t GD_COUNT = 11;


    static constexpr float  COEF_CN[] ={

        -4.98843114573573548651E-8, 9.50428062829859605134E-6,
               -6.45191435683965050962E-4, 1.88843319396703850064E-2,
               -2.05525900955013891793E-1, 9.99999999999999998822E-1

    };
    
    static constexpr float  COEF_CD[] ={

        3.99982968972495980367E-12, 9.15439215774657478799E-10,
               1.25001862479598821474E-7, 1.22262789024179030997E-5,
               8.68029542941784300606E-4, 4.12142090722199792936E-2,
               1.00000000000000000118E0

    };

    static constexpr float  COEF_SN[] ={

        -2.99181919401019853726E3, 7.08840045257738576863E5,
               -6.29741486205862506537E7, 2.54890880573376359104E9,
               -4.42979518059697779103E10, 3.18016297876567817986E11


    };
    static constexpr float  COEF_SD[] ={

               1.0,
               2.81376268889994315696E2,
               4.55847810806532581675E4,
               5.17343888770096400730E6,
               4.19320245898111231129E8,
               2.24411795645340920940E10,
               6.07366389490084639049E11
    };
    static constexpr float COEF_FN[] = {4.21543555043677546506E-1, 1.43407919780758885261E-1,
               1.15220955073585758835E-2, 3.45017939782574027900E-4,
               4.63613749287867322088E-6, 3.05568983790257605827E-8,
               1.02304514164907233465E-10, 1.72010743268161828879E-13,
               1.34283276233062758925E-16, 3.76329711269987889006E-20};

    static constexpr float COEF_FD[] = {1.0,
               7.51586398353378947175E-1,
               1.16888925859191382142E-1,
               6.44051526508858611005E-3,
               1.55934409164153020873E-4,
               1.84627567348930545870E-6,
               1.12699224763999035261E-8,
               3.60140029589371370404E-11,
               5.88754533621578410010E-14,
               4.52001434074129701496E-17,
               1.25443237090011264384E-20};

    static constexpr float COEF_GN[] = {5.04442073643383265887E-1, 1.97102833525523411709E-1,
               1.87648584092575249293E-2, 6.84079380915393090172E-4,
               1.15138826111884280931E-5, 9.82852443688422223854E-8,
               4.45344415861750144738E-10, 1.08268041139020870318E-12,
               1.37555460633261799868E-15, 8.36354435630677421531E-19,
               1.86958710162783235106E-22};
    static constexpr float COEF_GD[] = {1.0,
               1.47495759925128324529E0,
               3.37748989120019970451E-1,
               2.53603741420338795122E-2,
               8.14679107184306179049E-4,
               1.27545075667729118702E-5,
               1.04314589657571990585E-7,
               4.60680728146520428211E-10,
               1.10273215066240270757E-12,
               1.38796531259578871258E-15,
               8.39158816283118707363E-19,
               1.86958710162783236342E-22};
public:
    __aicore__ inline KernelFresnelCos() {}

    // template <bool IsExistBigCore>
    __aicore__ inline void Init(GM_ADDR x,
                                GM_ADDR y,
                                uint32_t ALIGN_NUM, uint32_t block_size, uint32_t core_size, uint32_t core_remain,
            TPipe * pipe)
    {
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->blockLength = core_size + (GetBlockNum() == GetBlockIdx() + 1 ? core_remain : 0);
        this->tileLength = block_size;
        this->ALIGN_NUM = ALIGN_NUM;
        this->blockLength = this->blockLength + (this->blockLength % ALIGN_NUM ? ALIGN_NUM - this->blockLength % ALIGN_NUM : 0);
        auto startPointer = core_size * GetBlockIdx();
        auto bufferlength = this->blockLength;

        xGm.SetGlobalBuffer((__gm__ T*)x + startPointer, bufferlength);
        yGm.SetGlobalBuffer((__gm__ T*)y + startPointer, bufferlength);

        this->tileNum = this->blockLength / this->tileLength + (this->blockLength % this->tileLength > 0);
        pipe->InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(T));
        pipe->InitBuffer(outQueueY, BUFFER_NUM,this->tileLength * sizeof(T));

        pipe->InitBuffer(B_w1, this->tileLength * sizeof(float));
        pipe->InitBuffer(B_w2, this->tileLength * sizeof(float));
        pipe->InitBuffer(tmp1, this->tileLength * sizeof(float));
        pipe->InitBuffer(tmp2, this->tileLength * sizeof(float));
        pipe->InitBuffer(tmp3, this->tileLength * sizeof(float));
        pipe->InitBuffer(tmp4, this->tileLength * sizeof(float));

        pipe->InitBuffer(tmp5, this->tileLength * sizeof(float));

        pipe->InitBuffer(tmp6, this->tileLength * sizeof(float));
        pipe->InitBuffer(tmp7, this->tileLength * sizeof(float));


        pipe->InitBuffer(B_bits1, this->tileLength * sizeof(uint8_t));
        pipe->InitBuffer(B_bits2, this->tileLength * sizeof(uint8_t));
        if constexpr (!std::is_same_v<T, float>) {
            pipe->InitBuffer(xBuf, this->tileLength * sizeof(float));
            pipe->InitBuffer(yBuf, this->tileLength * sizeof(float));
        }
        strategy.InitBufImpl(pipe, this->tileLength);


    }

    __aicore__ inline void Process()
    {
          int32_t loopCount = this->tileNum;
            for (int32_t i = 0; i < loopCount-1; i++) {
                CopyIn(i, this->tileLength);
                Compute(i, this->tileLength);
                CopyOut(i, this->tileLength);
            }
            uint32_t length = this->blockLength - this->tileLength * (loopCount - 1);
            CopyIn(loopCount - 1, length);
            Compute(loopCount - 1, length);
            CopyOut(loopCount - 1, length );
           
    }

private:
    __aicore__ inline void CopyIn(int32_t progress,uint32_t length)
    {
        AscendC::LocalTensor<T> xLocal = inQueueX.AllocTensor<T>();
        AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], length);
        inQueueX.EnQue(xLocal);
    }

    __aicore__ inline void Compute(int32_t progress,uint32_t length)
    {
        AscendC::LocalTensor<float> xLocal = PreDeQueCastX(length);
        AscendC::LocalTensor<float> yLocal = PreAllocateY();

        AscendC::LocalTensor<float> p1 = tmp1.Get<float>();

        AscendC::LocalTensor<float> p2 = tmp2.Get<float>();

        AscendC::LocalTensor<float> p3 = tmp3.Get<float>();

        AscendC::LocalTensor<float> p4 = tmp4.Get<float>();

        AscendC::LocalTensor<float> p5 = tmp5.Get<float>();


        AscendC::LocalTensor<float> p6 = tmp6.Get<float>();
        AscendC::LocalTensor<float> p7 = tmp7.Get<float>();

        AscendC::LocalTensor<uint8_t> bits1 = B_bits1.Get<uint8_t>();

        AscendC::LocalTensor<uint8_t> bits2 = B_bits2.Get<uint8_t>();

        AscendC::LocalTensor<float> w1 = B_w1.Get<float>();
        AscendC::LocalTensor<float> w2 = B_w2.Get<float>();


        //Duplicate(p1, static_cast<float>(0.0), length);
        Duplicate(p5, static_cast<float>(1.0), length);
        //Duplicate(p3, static_cast<float>(-1.0), length); 

        CompareScalar(bits1, xLocal,float(0.0), CMPMODE::GE, length);
        Select(p1, bits1, p5, float(-1.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, length); //p1不能动
   
        Abs(xLocal,xLocal,length);

        Mul(p6,xLocal,xLocal,length); //x^2

        Mul(p7,p6,p6,length);  //x^4

        //Duplicate(p2, float(THRESHOLD_B), length);
        CompareScalar(bits1, xLocal, float(THRESHOLD_B), AscendC::CMPMODE::GT, length);//大于THRESHOLD_B的,bit为1,小于等于THRESHOLD_B的,bit为0
        Duplicate(p3,float(0.50),length);
        Select(w1,bits1, p3, float(0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, length); //大于THRESHOLD_B的，选0,5,小于THRESHOLD_B的选0

        //Duplicate(p2,float(Sqrt_A),length);


        CompareScalar(bits2, xLocal,float(Sqrt_A), AscendC::CMPMODE::LT, length); //x^2小于THRESHOLD_A的,bit为1,大于等于THRESHOLD_A的,bit为0

        //输入是x^4

     

        // Mul(p3,xLocal,xLocal,length);

        // Mul(p3,p3,p3,length);


        _polevl(w2,p7,COEF_CN,CN_COUNT,length);
        _polevl(p2,p7,COEF_CD,CD_COUNT,length);

        Mul(w2,xLocal,w2,length);

        Div(p2,w2,p2,length);

        Select(w2,bits2, p2, float(0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, length);

        Or(w1.ReinterpretCast<uint16_t>(),w1.ReinterpretCast<uint16_t>(),w2.ReinterpretCast<uint16_t>(),length*2);//小于THRESHOLD_A的bit跟大于THRESHOLD_Bbit均为1，其余为0,最终结果

        Or(bits1.ReinterpretCast<uint16_t>(),bits1.ReinterpretCast<uint16_t>(),bits2.ReinterpretCast<uint16_t>(),length/2);

        Not(bits2.ReinterpretCast<uint16_t>(),bits1.ReinterpretCast<uint16_t>(),length/2); //a<=x<=b

        //现在可以用的buffer,p2,p3,w2,p4

        //Mul(p3,xLocal,xLocal,length);


        Muls(p3,p6,static_cast<float>(THRESHOLD_PI),length); // x^2 *pi

        //Duplicate(w2,static_cast<float>(1.0),length); // 1

        Div(p3,p5,p3,length);

        //Adds(p2,p3,static_cast<float>(0.0),length);

        Mul(p3,p3,p3,length);// (1/(x^2*pi))^2  可能有问题  暂时不能动

        
        _polevl(p2,p3,COEF_FN,FN_COUNT,length);//FN

        _polevl(w2,p3,COEF_FD,FD_COUNT,length);//FD

        Div(p2,p2,w2,length); // fn /fd

        Mul(p2,p3,p2,length); //data_f *(fn/fd)

        //Duplicate(w2,static_cast<float>(1.0),length);

        Sub(w2,p5,p2,length); //最终的data_f =w2

        _polevl(p2,p3,COEF_GN,GN_COUNT,length); //GN

        _polevl(p4,p3,COEF_GD,GD_COUNT,length); //GD

     

        Div(p3,p6,p7,length);

        Muls(p3,p3,static_cast<float>(ReciprocalPI),length); //data_rec

        Mul(p3,p2,p3,length); // data_rec * gn =data_g

        Div(p3,p3,p4,length); //data_g /gd  最终的data_g =p3  上述操作结果没出问题

 


        Muls(p4,p6,static_cast<float>(xishu),length); //data_z

  

        strategy.ComputeImplCos(p4,p2,length); //cosx  data_c = p2   p2 =cosx

        //Duplicate(yLocal,static_cast<float>(4.0),length);

        Muls(p7,p6,float(0.25),length);

        Cast(p7,p7,RoundMode::CAST_TRUNC,length);
        Muls(p7,p7,float(4.0),length);

        Sub(yLocal,p6,p7,length);   // mod4 = x2 % 4.0  //5步  yLocal不动了

        //Duplicate(p4,static_cast<float>(2.0),length);


        CompareScalar(bits1, yLocal, (float)(2.0), CMPMODE::LE, length);

        //Duplicate(p4,static_cast<float>(1.0),length);

        Select(p7, bits1, p5, float(-1.0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, length); //p7不能动

        //p7:    <2 ,p7=1 , >2,p7=-1


        Mul(p4,p2,p2,length); //cosx^2

        //Duplicate(yLocal,static_cast<float>(1.0),length);

        Sub(p4,p5,p4,length);

        Sqrt(yLocal,p4,length);

        Mul(p7,p7,yLocal,length);  //14步


        //strategy.ComputeImplSin(yLocal,p7,length); //sin  data_s = p5


        Muls(yLocal,xLocal,static_cast<float>(THRESHOLD_PI),length); // data_y

        //Duplicate(p4,static_cast<float>(1.0),length);

        Div(yLocal,p5,yLocal,length); //data_y


        Mul(p4,p7,w2,length);

        Mul(p7,p3,p2,length);

        Sub(p4,p4,p7,length);

        Mul(p4,yLocal,p4,length);

        Adds(p4,p4,static_cast<float>(0.50),length);

        Select(yLocal,bits2, p4, float(0), AscendC::SELMODE::VSEL_TENSOR_SCALAR_MODE, length);

        Or(yLocal.ReinterpretCast<uint16_t>(), yLocal.ReinterpretCast<uint16_t>(), w1.ReinterpretCast<uint16_t>(), length * 2);

        Mul(yLocal,p1,yLocal,length);

        PostReleaseCastEnQue(xLocal, yLocal,length);
    }
    
    //这里的w其实本质是就是输入的x,这里的coef本质上是对应不同的COEF参数
    __aicore__ inline void _polevl(AscendC::LocalTensor<float> &dst, AscendC::LocalTensor<float> &w, const float *coef,uint32_t N, uint32_t length) {


        AscendC::Duplicate(dst, coef[0], length);  
        for (int i = 1; i <= N; ++i) {


            AscendC::Mul(dst, dst, w, length);
            AscendC::Adds(dst, dst, coef[i], length); //这一步是没问题的

            // Duplicate( jiashu,(float)(coef[i]),length);

            // AscendC::FusedMulAdd(dst,w, jiashu, length);
        }

       
        
    }

    __aicore__ inline void CopyOut(int32_t progress,uint32_t length)
    {
        AscendC::LocalTensor<T> yLocal = outQueueY.DeQue<T>();
        AscendC::DataCopy(yGm[progress * this->tileLength], yLocal, length);
        outQueueY.FreeTensor(yLocal);
    }

    __aicore__ inline AscendC::LocalTensor<float> PreDeQueCastX(uint32_t length)
    {
        if constexpr (std::is_same_v<T, float>) {
            AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
            return xLocal;
        } else {
            AscendC::LocalTensor<float> xLocal = xBuf.Get<float>();
            AscendC::LocalTensor<T> xOrigin = inQueueX.DeQue<T>();
            AscendC::Cast(xLocal, xOrigin, AscendC::RoundMode::CAST_NONE, length);
            inQueueX.FreeTensor(xOrigin);
            return xLocal;
        }
    }



    __aicore__ inline AscendC::LocalTensor<float> PreAllocateY()
    {
        if constexpr (std::is_same_v<T, float>) {
            AscendC::LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
            return yLocal;
        } else {
            AscendC::LocalTensor<float> yLocal = yBuf.Get<float>();
            return yLocal;
        }
    }

    __aicore__ inline void PostReleaseCastEnQue(AscendC::LocalTensor<float>& xLocal,
                                                AscendC::LocalTensor<float>& yLocal,uint32_t length)
    {
        if constexpr (std::is_same_v<T, float>) {
            outQueueY.EnQue(yLocal);
            inQueueX.FreeTensor(xLocal);
        } else {
            AscendC::LocalTensor<T> yTarget = outQueueY.AllocTensor<T>();
            AscendC::Cast(yTarget, yLocal, AscendC::RoundMode::CAST_RINT, length);
            outQueueY.EnQue(yTarget);
        }
    }

private:
    AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX;
    AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    AscendC::TBuf<AscendC::QuePosition::VECCALC> B_w1,B_w2,xBuf, yBuf,tmp1,tmp2,tmp3,tmp4,tmp5,tmp6,tmp7,B_bits1,B_bits2;
    AscendC::GlobalTensor<T> xGm;
    AscendC::GlobalTensor<T> yGm;

    uint32_t blockLength;
    uint32_t tileNum;
    uint32_t tileLength;
    uint32_t ALIGN_NUM;

    ComputeStrategy strategy;
};

class HighPerfStrategyCos
{
public:
    __aicore__ inline HighPerfStrategyCos() {}

    __aicore__ inline void InitBufImpl(AscendC::TPipe* pipe, uint32_t ubPartDataNum)
    {
        pipe->InitBuffer(tmpBuf1, ubPartDataNum * sizeof(float));
        pipe->InitBuffer(tmpBuf2, ubPartDataNum * sizeof(float));
        pipe->InitBuffer(tmpBuf3, ubPartDataNum * sizeof(float));
        pipe->InitBuffer(tmpBuf4, ubPartDataNum * sizeof(float));
    }

    __aicore__ inline void ComputeImplCos(AscendC::LocalTensor<float>& xLocal,
                                       AscendC::LocalTensor<float>& yLocal,
                                       uint32_t processDataNum);
    
    // __aicore__ inline void ComputeImplSin(AscendC::LocalTensor<float>& xLocal,
    //                                    AscendC::LocalTensor<float>& yLocal,
    //                                    uint32_t processDataNum);

private:
    AscendC::TBuf<AscendC::QuePosition::VECCALC> tmpBuf1, tmpBuf2, tmpBuf3, tmpBuf4;
};

constexpr float PI_FOR_X_TODIV = 0.3183098733425140380859375f;

constexpr float PI_DOWN = 1.57079637050628662109375f;
constexpr float PI_RESDOWN_ADDS_NEG = -0.00000004371139000189375f;

constexpr float COS_RES_MULIT_SCA = 2.604926501e-6f;
constexpr float COS_RES_ADDICT_UP = -0.0001980894471f;
constexpr float COS_2ADDS = 0.008333049340f;
constexpr float COS_3ADDS = -0.1666665792f;

// constexpr float SIN_RES_MULIT_SCA = 2.604926501e-6f;
// constexpr float SIN_RES_ADDICT_UP = -0.0001980894471f;
// constexpr float SIN_2ADDS = 0.008333049340f;
// constexpr float SIN_3ADDS = -0.1666665792f;

constexpr float pi_0 = 3.14160156f;
constexpr float pi_1 = -8.9071691e-06f;
constexpr float pi_2 = -1.74122761e-09f;
constexpr float pi_3 = 1.24467439e-13f;


__aicore__ inline void HighPerfStrategyCos::ComputeImplCos(AscendC::LocalTensor<float>& xLocal,
                                                     AscendC::LocalTensor<float>& yLocal,
                                                     uint32_t processDataNum)
{
    AscendC::LocalTensor<float> tmpTensor1 = tmpBuf1.Get<float>();
    AscendC::LocalTensor<float> tmpTensor2 = tmpBuf2.Get<float>();
    AscendC::LocalTensor<float> tmpTensor3 = tmpBuf3.Get<float>();
    AscendC::LocalTensor<float> tmpTensor4 = tmpBuf4.Get<float>();

    const AscendC::LocalTensor<float>& input_x = xLocal;
    const AscendC::LocalTensor<float>& x_vmul = tmpTensor1;
    const AscendC::LocalTensor<float>& x_vmul1 = tmpTensor2;
    const AscendC::LocalTensor<float>& x_vmul0 = yLocal;
    const AscendC::LocalTensor<float>& round_pi_div = tmpTensor1;
    const AscendC::LocalTensor<float>& round_pi_div0 = tmpTensor3;
    const AscendC::LocalTensor<float>& round_pi_div0_1 = tmpTensor2;
    const AscendC::LocalTensor<float>& round_pi_div1 = yLocal;
    const AscendC::LocalTensor<float>& fix = tmpTensor4;
    const AscendC::LocalTensor<float>& x_fixed = tmpTensor3;
    const AscendC::LocalTensor<float>& fix_1 = tmpTensor4;
    const AscendC::LocalTensor<float>& x_fixed_1 = xLocal;
    const AscendC::LocalTensor<float>& fix_2 = tmpTensor4;
    const AscendC::LocalTensor<float>& x_fixed_2 = tmpTensor3;
    const AscendC::LocalTensor<float>& x_fixed_3 = xLocal;
    const AscendC::LocalTensor<float>& fix_3 = tmpTensor4;
    const AscendC::LocalTensor<float>& x_fixed_4 = tmpTensor3;
    const AscendC::LocalTensor<float>& fix_4 = tmpTensor4;
    const AscendC::LocalTensor<float>& x_fixed_5 = xLocal;
    const AscendC::LocalTensor<float>& fix_5 = tmpTensor4;
    const AscendC::LocalTensor<float>& x_fixed_6 = tmpTensor3;
    const AscendC::LocalTensor<float>& fix_6 = xLocal;
    const AscendC::LocalTensor<float>& x_fixed_7 = tmpTensor2;
    const AscendC::LocalTensor<float>& fix_7 = xLocal;
    const AscendC::LocalTensor<float>& x_fixed_8 = tmpTensor3;
    const AscendC::LocalTensor<float>& x_fixed_9 = yLocal;
    const AscendC::LocalTensor<float>& x_pow = tmpTensor2;
    const AscendC::LocalTensor<float>& kover2 = xLocal;
    const AscendC::LocalTensor<float>& kover2floor = tmpTensor3;
    const AscendC::LocalTensor<float>& kover2floorm4 = xLocal;
    const AscendC::LocalTensor<float>& k2 = tmpTensor3;
    const AscendC::LocalTensor<float>& sign = tmpTensor4;
    const AscendC::LocalTensor<float>& sign_1 = tmpTensor1;
    const AscendC::LocalTensor<float>& res_up = tmpTensor3;
    const AscendC::LocalTensor<float>& res_up_1 = xLocal;
    const AscendC::LocalTensor<float>& res_up_2 = tmpTensor3;
    const AscendC::LocalTensor<float>& res_up_3 = xLocal;
    const AscendC::LocalTensor<float>& res_up_4 = tmpTensor3;
    const AscendC::LocalTensor<float>& res_up_5 = xLocal;
    const AscendC::LocalTensor<float>& res_up_6 = tmpTensor3;
    const AscendC::LocalTensor<float>& res_up_7 = tmpTensor2;
    const AscendC::LocalTensor<float>& res_up_8 = xLocal;
    const AscendC::LocalTensor<float>& res_sign = yLocal;
    const AscendC::LocalTensor<float>& res_mins = tmpTensor1;
    const AscendC::LocalTensor<float>& res_maxs = yLocal;

    AscendC::Muls(x_vmul, input_x, PI_FOR_X_TODIV, processDataNum);
    AscendC::Adds(x_vmul1, x_vmul, 0.5f, processDataNum);
    AscendC::Muls(x_vmul0, x_vmul, 1.0f / 2048.0f, processDataNum);
    AscendC::Cast(round_pi_div, x_vmul1, AscendC::RoundMode::CAST_ROUND, processDataNum);
    AscendC::Cast(round_pi_div0, x_vmul0, AscendC::RoundMode::CAST_ROUND, processDataNum);
    AscendC::Muls(round_pi_div0_1, round_pi_div0, 2048.0f, processDataNum);
    AscendC::Sub(round_pi_div1, round_pi_div, round_pi_div0_1, processDataNum);

    AscendC::Muls(fix, round_pi_div0_1, pi_0, processDataNum);
    AscendC::Sub(x_fixed, input_x, fix, processDataNum);
    AscendC::Muls(fix_1, round_pi_div1, pi_0, processDataNum);
    AscendC::Sub(x_fixed_1, x_fixed, fix_1, processDataNum);
    AscendC::Muls(fix_2, round_pi_div0_1, pi_1, processDataNum);
    AscendC::Sub(x_fixed_2, x_fixed_1, fix_2, processDataNum);

    AscendC::Adds(x_fixed_3, x_fixed_2, PI_DOWN, processDataNum);

    AscendC::Muls(fix_3, round_pi_div1, pi_1, processDataNum);
    AscendC::Sub(x_fixed_4, x_fixed_3, fix_3, processDataNum);
    AscendC::Muls(fix_4, round_pi_div0_1, pi_2, processDataNum);
    AscendC::Sub(x_fixed_5, x_fixed_4, fix_4, processDataNum);
    AscendC::Muls(fix_5, round_pi_div1, pi_2, processDataNum);
    AscendC::Sub(x_fixed_6, x_fixed_5, fix_5, processDataNum);
    AscendC::Muls(fix_6, round_pi_div0_1, pi_3, processDataNum);
    AscendC::Sub(x_fixed_7, x_fixed_6, fix_6, processDataNum);
    AscendC::Muls(fix_7, round_pi_div1, pi_3, processDataNum);
    AscendC::Sub(x_fixed_8, x_fixed_7, fix_7, processDataNum);
    AscendC::Adds(x_fixed_9, x_fixed_8, PI_RESDOWN_ADDS_NEG, processDataNum);

    AscendC::Mul(x_pow, x_fixed_9, x_fixed_9, processDataNum);
    AscendC::Muls(kover2, round_pi_div, 0.5f, processDataNum);
    AscendC::Cast(kover2floor, kover2, AscendC::RoundMode::CAST_FLOOR, processDataNum);
    AscendC::Muls(kover2floorm4, kover2floor, 4.0f, processDataNum);
    AscendC::Muls(k2, round_pi_div, -2.0f, processDataNum);
    AscendC::Add(sign, kover2floorm4, k2, processDataNum);
    AscendC::Adds(sign_1, sign, 1.0f, processDataNum);

    AscendC::Muls(res_up, x_pow, COS_RES_MULIT_SCA, processDataNum);
    AscendC::Adds(res_up_1, res_up, COS_RES_ADDICT_UP, processDataNum);
    AscendC::Mul(res_up_2, res_up_1, x_pow, processDataNum);
    AscendC::Adds(res_up_3, res_up_2, COS_2ADDS, processDataNum);
    AscendC::Mul(res_up_4, res_up_3, x_pow, processDataNum);
    AscendC::Adds(res_up_5, res_up_4, COS_3ADDS, processDataNum);
    AscendC::Mul(res_up_6, res_up_5, x_pow, processDataNum);
    AscendC::Adds(res_up_7, res_up_6, 1.0f, processDataNum);
    AscendC::Mul(res_up_8, res_up_7, x_fixed_9, processDataNum);
    AscendC::Mul(res_sign, res_up_8, sign_1, processDataNum);

    // Ensure result is between -1 and 1
    AscendC::Mins(res_mins, res_sign, 1.0f, processDataNum);
    AscendC::Maxs(res_maxs, res_mins, -1.0f, processDataNum); //44步
}

// __aicore__ inline void HighPerfStrategyCos::ComputeImplCos(AscendC::LocalTensor<float>& xLocal,
//                                                      AscendC::LocalTensor<float>& yLocal,
//                                                      uint32_t processDataNum)
// {
//     AscendC::LocalTensor<float> tmpTensor1 = tmpBuf1.Get<float>();
//     AscendC::LocalTensor<float> tmpTensor2 = tmpBuf2.Get<float>();
//     AscendC::LocalTensor<float> tmpTensor3 = tmpBuf3.Get<float>();
//     AscendC::LocalTensor<float> tmpTensor4 = tmpBuf4.Get<float>();

//     const AscendC::LocalTensor<float>& input_x = xLocal;
//     const AscendC::LocalTensor<float>& x_vmul = tmpTensor1;
//     const AscendC::LocalTensor<float>& x_vmul1 = tmpTensor2;
//     const AscendC::LocalTensor<float>& x_vmul0 = yLocal;
//     const AscendC::LocalTensor<float>& round_pi_div = tmpTensor1;
//     const AscendC::LocalTensor<float>& round_pi_div0 = tmpTensor3;
//     const AscendC::LocalTensor<float>& round_pi_div0_1 = tmpTensor2;
//     const AscendC::LocalTensor<float>& round_pi_div1 = yLocal;
//     const AscendC::LocalTensor<float>& fix = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed = tmpTensor3;
//     const AscendC::LocalTensor<float>& fix_1 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_1 = xLocal;
//     const AscendC::LocalTensor<float>& fix_2 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_2 = tmpTensor3;
//     const AscendC::LocalTensor<float>& x_fixed_3 = xLocal;
//     const AscendC::LocalTensor<float>& fix_3 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_4 = tmpTensor3;
//     const AscendC::LocalTensor<float>& fix_4 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_5 = xLocal;
//     const AscendC::LocalTensor<float>& fix_5 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_6 = tmpTensor3;
//     const AscendC::LocalTensor<float>& fix_6 = xLocal;
//     const AscendC::LocalTensor<float>& x_fixed_7 = tmpTensor2;
//     const AscendC::LocalTensor<float>& fix_7 = xLocal;
//     const AscendC::LocalTensor<float>& x_fixed_8 = tmpTensor3;
//     const AscendC::LocalTensor<float>& x_fixed_9 = yLocal;
//     const AscendC::LocalTensor<float>& x_pow = tmpTensor2;
//     const AscendC::LocalTensor<float>& kover2 = xLocal;
//     const AscendC::LocalTensor<float>& kover2floor = tmpTensor3;
//     const AscendC::LocalTensor<float>& kover2floorm4 = xLocal;
//     const AscendC::LocalTensor<float>& k2 = tmpTensor3;
//     const AscendC::LocalTensor<float>& sign = tmpTensor4;
//     const AscendC::LocalTensor<float>& sign_1 = tmpTensor1;
//     const AscendC::LocalTensor<float>& res_up = tmpTensor3;
//     const AscendC::LocalTensor<float>& res_up_1 = xLocal;
//     const AscendC::LocalTensor<float>& res_up_2 = tmpTensor3;
//     const AscendC::LocalTensor<float>& res_up_3 = xLocal;
//     const AscendC::LocalTensor<float>& res_up_4 = tmpTensor3;
//     const AscendC::LocalTensor<float>& res_up_5 = xLocal;
//     const AscendC::LocalTensor<float>& res_up_6 = tmpTensor3;
//     const AscendC::LocalTensor<float>& res_up_7 = tmpTensor2;
//     const AscendC::LocalTensor<float>& res_up_8 = xLocal;
//     const AscendC::LocalTensor<float>& res_sign = yLocal;
//     const AscendC::LocalTensor<float>& res_mins = tmpTensor1;
//     const AscendC::LocalTensor<float>& res_maxs = yLocal;

//     AscendC::Muls(x_vmul, input_x, PI_FOR_X_TODIV, processDataNum);
//     AscendC::Adds(x_vmul1, x_vmul, 0.5f, processDataNum);
//     AscendC::Muls(x_vmul0, x_vmul, 1.0f / 2048.0f, processDataNum);
//     AscendC::Cast(round_pi_div, x_vmul1, AscendC::RoundMode::CAST_ROUND, processDataNum);
//     AscendC::Cast(round_pi_div0, x_vmul0, AscendC::RoundMode::CAST_ROUND, processDataNum);
//     AscendC::Muls(round_pi_div0_1, round_pi_div0, 2048.0f, processDataNum);
//     AscendC::Sub(round_pi_div1, round_pi_div, round_pi_div0_1, processDataNum); //7步

//     AscendC::Muls(fix, round_pi_div0_1, pi_0, processDataNum);
//     AscendC::Sub(x_fixed, input_x, fix, processDataNum);
//     AscendC::Muls(fix_1, round_pi_div1, pi_0, processDataNum);
//     AscendC::Sub(x_fixed_1, x_fixed, fix_1, processDataNum);
//     AscendC::Muls(fix_2, round_pi_div0_1, pi_1, processDataNum);
//     AscendC::Sub(x_fixed_2, x_fixed_1, fix_2, processDataNum);

//     AscendC::Adds(x_fixed_3, x_fixed_2, PI_DOWN, processDataNum);

//     AscendC::Muls(fix_3, round_pi_div1, pi_1, processDataNum);
//     AscendC::Sub(x_fixed_4, x_fixed_3, fix_3, processDataNum);
//     AscendC::Muls(fix_4, round_pi_div0_1, pi_2, processDataNum);
//     AscendC::Sub(x_fixed_5, x_fixed_4, fix_4, processDataNum);
//     AscendC::Muls(fix_5, round_pi_div1, pi_2, processDataNum);
//     AscendC::Sub(x_fixed_6, x_fixed_5, fix_5, processDataNum);
//     AscendC::Muls(fix_6, round_pi_div0_1, pi_3, processDataNum);
//     AscendC::Sub(x_fixed_7, x_fixed_6, fix_6, processDataNum);
//     AscendC::Muls(fix_7, round_pi_div1, pi_3, processDataNum);
//     AscendC::Sub(x_fixed_8, x_fixed_7, fix_7, processDataNum);
//     AscendC::Adds(x_fixed_9, x_fixed_8, PI_RESDOWN_ADDS_NEG, processDataNum);

//     AscendC::Mul(x_pow, x_fixed_9, x_fixed_9, processDataNum);  //19步


   

//     AscendC::Muls(kover2, round_pi_div, 0.5f, processDataNum);
//     AscendC::Cast(kover2floor, kover2, AscendC::RoundMode::CAST_FLOOR, processDataNum);
//     AscendC::Muls(kover2floorm4, kover2floor, 4.0f, processDataNum);
//     AscendC::Muls(k2, round_pi_div, -2.0f, processDataNum);
//     AscendC::Add(sign, kover2floorm4, k2, processDataNum);
//     AscendC::Adds(sign_1, sign, 1.0f, processDataNum);

   
//     AscendC::Muls(res_up, x_pow, COS_2ADDS, processDataNum);           // res_up = COS_2ADDS * x^2
//     AscendC::Adds(res_up_1, res_up, COS_3ADDS, processDataNum);        // res_up_1 = COS_2ADDS*x^2 + COS_3ADDS
//     AscendC::Mul(res_up_2, res_up_1, x_pow, processDataNum);           // res_up_2 = (..) * x^2 → COS_2ADDS*x^4 + COS_3ADDS*x^2
//     AscendC::Adds(res_up_3, res_up_2, 1.0f, processDataNum);           // res_up_3 = P(x^2) = 1 + COS_3ADDS*x^2 + COS_2ADDS*x^4

//     AscendC::Mul(res_up_8, res_up_3, x_fixed_9, processDataNum);       // res_up_8 = x_fixed_9 * P(x^2) ≈ sin(x')
//     AscendC::Mul(res_sign, res_up_8, sign_1, processDataNum);          // res_sign = sin(x') * sign → cos(x)

//     // Ensure result is between -1 and 1
//     AscendC::Mins(res_mins, res_sign, 1.0f, processDataNum);
//     AscendC::Maxs(res_maxs, res_mins, -1.0f, processDataNum);  //14步    //40步
// }


// __aicore__ inline void HighPerfStrategyCos::ComputeImplSin(AscendC::LocalTensor<float>& xLocal,
//                                                      AscendC::LocalTensor<float>& yLocal,
//                                                      uint32_t processDataNum)
// {
//     AscendC::LocalTensor<float> tmpTensor1 = tmpBuf1.Get<float>();
//     AscendC::LocalTensor<float> tmpTensor2 = tmpBuf2.Get<float>();
//     AscendC::LocalTensor<float> tmpTensor3 = tmpBuf3.Get<float>();
//     AscendC::LocalTensor<float> tmpTensor4 = tmpBuf4.Get<float>();

//     const AscendC::LocalTensor<float>& input_x = xLocal;
//     const AscendC::LocalTensor<float>& x_vmul = tmpTensor2;
//     const AscendC::LocalTensor<float>& x_vmul0 = yLocal;
//     const AscendC::LocalTensor<float>& round_pi_div = tmpTensor1;
//     const AscendC::LocalTensor<float>& round_pi_div0 = tmpTensor3;
//     const AscendC::LocalTensor<float>& round_pi_div0_1 = tmpTensor2;
//     const AscendC::LocalTensor<float>& round_pi_div1 = yLocal;
//     const AscendC::LocalTensor<float>& fix = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed = tmpTensor3;
//     const AscendC::LocalTensor<float>& fix_1 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_1 = xLocal;
//     const AscendC::LocalTensor<float>& fix_2 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_2 = tmpTensor3;
//     const AscendC::LocalTensor<float>& fix_3 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_3 = xLocal;
//     const AscendC::LocalTensor<float>& fix_4 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_4 = tmpTensor3;
//     const AscendC::LocalTensor<float>& fix_5 = tmpTensor4;
//     const AscendC::LocalTensor<float>& x_fixed_5 = xLocal;
//     const AscendC::LocalTensor<float>& fix_6 = tmpTensor3;
//     const AscendC::LocalTensor<float>& x_fixed_6 = tmpTensor2;
//     const AscendC::LocalTensor<float>& fix_7 = xLocal;
//     const AscendC::LocalTensor<float>& x_fixed_7 = yLocal;
//     const AscendC::LocalTensor<float>& x_pow = tmpTensor2;
//     const AscendC::LocalTensor<float>& kover2 = xLocal;
//     const AscendC::LocalTensor<float>& kover2floor = tmpTensor3;
//     const AscendC::LocalTensor<float>& kover2floorm4 = xLocal;
//     const AscendC::LocalTensor<float>& k2 = tmpTensor3;
//     const AscendC::LocalTensor<float>& sign = tmpTensor4;
//     const AscendC::LocalTensor<float>& sign_1 = tmpTensor1;
//     const AscendC::LocalTensor<float>& res_up = tmpTensor3;
//     const AscendC::LocalTensor<float>& res_up_1 = xLocal;
//     const AscendC::LocalTensor<float>& res_up_2 = tmpTensor3;
//     const AscendC::LocalTensor<float>& res_up_3 = xLocal;
//     const AscendC::LocalTensor<float>& res_up_4 = tmpTensor3;
//     const AscendC::LocalTensor<float>& res_up_5 = xLocal;
//     const AscendC::LocalTensor<float>& res_up_6 = tmpTensor2;
//     //const AscendC::LocalTensor<float>& res_up_6 = tmpTensor3;
//     // const AscendC::LocalTensor<float>& res_up_7 = tmpTensor2;
//     // const AscendC::LocalTensor<float>& res_up_8 = xLocal;
//     const AscendC::LocalTensor<float>& res_sign = yLocal;
//     const AscendC::LocalTensor<float>& res_mins = tmpTensor1;
//     const AscendC::LocalTensor<float>& res_maxs = yLocal;

//     AscendC::Muls(x_vmul, input_x, PI_FOR_X_TODIV, processDataNum);
//     AscendC::Muls(x_vmul0, x_vmul, 1.0f / 2048.0f, processDataNum);
//     AscendC::Cast(round_pi_div, x_vmul, AscendC::RoundMode::CAST_ROUND, processDataNum);
//     AscendC::Cast(round_pi_div0, x_vmul0, AscendC::RoundMode::CAST_ROUND, processDataNum);
//     AscendC::Muls(round_pi_div0_1, round_pi_div0, 2048.0f, processDataNum);
//     AscendC::Sub(round_pi_div1, round_pi_div, round_pi_div0_1, processDataNum);
//     AscendC::Muls(fix, round_pi_div0_1, pi_0, processDataNum);
//     AscendC::Sub(x_fixed, input_x, fix, processDataNum);
//     AscendC::Muls(fix_1, round_pi_div1, pi_0, processDataNum);
//     AscendC::Sub(x_fixed_1, x_fixed, fix_1, processDataNum);
//     AscendC::Muls(fix_2, round_pi_div0_1, pi_1, processDataNum);
//     AscendC::Sub(x_fixed_2, x_fixed_1, fix_2, processDataNum);
//     AscendC::Muls(fix_3, round_pi_div1, pi_1, processDataNum);
//     AscendC::Sub(x_fixed_3, x_fixed_2, fix_3, processDataNum);
//     AscendC::Muls(fix_4, round_pi_div0_1, pi_2, processDataNum);
//     AscendC::Sub(x_fixed_4, x_fixed_3, fix_4, processDataNum);
//     AscendC::Muls(fix_5, round_pi_div1, pi_2, processDataNum);
//     AscendC::Sub(x_fixed_5, x_fixed_4, fix_5, processDataNum);
//     AscendC::Muls(fix_6, round_pi_div0_1, pi_3, processDataNum);
//     AscendC::Sub(x_fixed_6, x_fixed_5, fix_6, processDataNum);
//     AscendC::Muls(fix_7, round_pi_div1, pi_3, processDataNum);
//     AscendC::Sub(x_fixed_7, x_fixed_6, fix_7, processDataNum);
//     AscendC::Mul(x_pow, x_fixed_7, x_fixed_7, processDataNum);   //23步 

//     //
//     AscendC::Muls(res_up, x_pow, SIN_RES_ADDICT_UP, processDataNum);     // z^2 * a7
//     AscendC::Adds(res_up_1, res_up, SIN_2ADDS, processDataNum);           // + a5
//     AscendC::Mul(res_up_2, res_up_1, x_pow, processDataNum);              // * z^2 → z^4 * a5 + z^6 * a7
//     AscendC::Adds(res_up_3, res_up_2, SIN_3ADDS, processDataNum);         // + a3
//     AscendC::Mul(res_up_4, res_up_3, x_pow, processDataNum);              // * z^2 → z^2 * a3 + ...
//     AscendC::Adds(res_up_5, res_up_4, 1.0f, processDataNum);               // + 1
//     AscendC::Mul(res_up_6, res_up_5, x_fixed_7, processDataNum);           // * z → sin(z) 近似

//     // Step 6: 符号修正（根据 k mod 4）
//     AscendC::Muls(kover2, round_pi_div, 0.5f, processDataNum);
//     AscendC::Cast(kover2floor, kover2, AscendC::RoundMode::CAST_FLOOR, processDataNum);
//     AscendC::Muls(kover2floorm4, kover2floor, 4.0f, processDataNum);
//     AscendC::Muls(k2, round_pi_div, -2.0f, processDataNum);
//     AscendC::Add(sign, kover2floorm4, k2, processDataNum);
//     AscendC::Adds(sign_1, sign, 1.0f, processDataNum);                     // sign adjustment

//     // 应用符号
//     AscendC::Mul(res_sign, res_up_6, sign_1, processDataNum);

//     // Step 7: 裁剪到 [-1, 1]（防止浮点误差溢出）
//     AscendC::Mins(res_mins, res_sign, 1.0f, processDataNum);
//     AscendC::Maxs(res_maxs, res_mins, -1.0f, processDataNum);  // 16步
    //

    // AscendC::Muls(kover2, round_pi_div, 0.5f, processDataNum);
    // AscendC::Cast(kover2floor, kover2, AscendC::RoundMode::CAST_FLOOR, processDataNum);
    // AscendC::Muls(kover2floorm4, kover2floor, 4.0f, processDataNum);
    // AscendC::Muls(k2, round_pi_div, -2.0f, processDataNum);
    // AscendC::Add(sign, kover2floorm4, k2, processDataNum);
    // AscendC::Adds(sign_1, sign, 1.0f, processDataNum);

    // AscendC::Muls(res_up, x_pow, SIN_RES_MULIT_SCA, processDataNum);
    // AscendC::Adds(res_up_1, res_up, SIN_RES_ADDICT_UP, processDataNum);
    // AscendC::Mul(res_up_2, res_up_1, x_pow, processDataNum);
    // AscendC::Adds(res_up_3, res_up_2, SIN_2ADDS, processDataNum);


    // AscendC::Mul(res_up_4, res_up_3, x_pow, processDataNum);
    // AscendC::Adds(res_up_5, res_up_4, SIN_3ADDS, processDataNum);
    // AscendC::Mul(res_up_6, res_up_5, x_pow, processDataNum);
    // AscendC::Adds(res_up_7, res_up_6, 1.0f, processDataNum);
    // AscendC::Mul(res_up_8, res_up_7, x_fixed_7, processDataNum);
    // AscendC::Mul(res_sign, res_up_8, sign_1, processDataNum);
    // AscendC::Mins(res_mins, res_sign, 1.0f, processDataNum);
    // AscendC::Maxs(res_maxs, res_mins, -1.0f, processDataNum);
//}




extern "C" __global__ __aicore__ void fresnel_cos(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(tiling_data, tiling);


    using ComputeStrategy = HighPerfStrategyCos;

    KernelFresnelCos<DTYPE_X, ComputeStrategy> op;
    AscendC::TPipe pipe;

    op.Init(x,y,tiling_data.ALIGN_NUM, tiling_data.block_size, tiling_data.core_size, tiling_data.core_remain,&pipe);

    op.Process();
}